/*
 * Projekt: GMU 2013 - Canny hranovy detektor OpenCL/CUDA
 * Autori: Pavel Nevrala - xnevra00 (xnevra00@stud.fit.vutbr.cz)
 *         Martin Jurak - xjurak02 (xjurak02@stud.fit.vutbr.cz)
 */
#include "NonMaxSupCL.h"


NonMaxSupCL::NonMaxSupCL(cl_uint maxWorkGroupP, cl_float *edgesDirectionP, cl_context *contextP, cl_command_queue *commandQueueP, cl_program *programP): 
	CannyPartsBase(maxWorkGroupP), edgesDirection(edgesDirectionP), context(contextP), commandQueue(commandQueueP), program(programP)
{
}


NonMaxSupCL::~NonMaxSupCL(void)
{
}


/*
 * Provede nastaveni OpenCl a kernelu pro hledani lokalnich maxim v obrazu 
 * @image - vstupni obraz
 * @width - sirka obrazu
 * @height - vyska obrazu
 * @return - doba vypoctu na GPU
 */
double NonMaxSupCL::run(cl_float *image, cl_uint width, cl_uint height)
{
	cl_int clErr;

	//vytvoreni bufferu
	// vstupni obraz
	cl_mem cl_in = clCreateBuffer(*context, CL_MEM_READ_ONLY, width * height * sizeof(cl_float), NULL, &clErr);
	clErr = clEnqueueWriteBuffer(*commandQueue, cl_in, CL_TRUE /*blocking write*/,	0, width * height * sizeof(cl_float), image, 0, NULL, NULL);
	// smer gradient jednotlivych pixelu
	cl_mem cl_edgesDirection = clCreateBuffer(*context, CL_MEM_READ_ONLY, width * height * sizeof(cl_float), NULL, &clErr);
	clErr = clEnqueueWriteBuffer(*commandQueue, cl_edgesDirection, CL_TRUE /*blocking write*/,	0, width * height * sizeof(cl_float), edgesDirection, 0, NULL, NULL);
	
	// vystupni obraz
	cl_mem cl_outSup = clCreateBuffer(*context, CL_MEM_WRITE_ONLY, width * height * sizeof(cl_float), NULL, &clErr);	


	//ziskani kernelu	
	cl_kernel kernelNonMaxSup = clCreateKernel(*program, "nonMaxSup", &clErr);
	errorController.CheckOpenCLError( clErr, "clCreateKernel kernelnonMaxSup" );	
	
// ========================================================================================================================================================================================
// ========================================================================================================================================================================================

	clErr = clSetKernelArg(kernelNonMaxSup, 0, sizeof(cl_mem), &cl_in);
	errorController.CheckOpenCLError(clErr, "clSetKernelArg. (inputImage)");
	clErr = clSetKernelArg(kernelNonMaxSup, 1, sizeof(cl_mem), &cl_edgesDirection);
	errorController.CheckOpenCLError(clErr, "clSetKernelArg. (outImage)");
	clErr = clSetKernelArg(kernelNonMaxSup, 2, sizeof(cl_mem), &cl_outSup);
	errorController.CheckOpenCLError(clErr, "clSetKernelArg. (outImage)");
	clErr = clSetKernelArg(kernelNonMaxSup, 3, sizeof(cl_uint), &width);
	errorController.CheckOpenCLError(clErr, "clSetKernelArg. (widthImage)");
	clErr = clSetKernelArg(kernelNonMaxSup, 4, sizeof(cl_uint), &height);
	errorController.CheckOpenCLError(clErr, "clSetKernelArg. (heightImage)");

	size_t GWS[2], LWS[2];
	GWS[0]=maxWorkGroup;
	GWS[1]=height;
	LWS[0]=maxWorkGroup;
	LWS[1]=1;

	cl_event event_NonMaxSup;

	clErr = clEnqueueNDRangeKernel(*commandQueue, kernelNonMaxSup, 2, NULL, GWS, LWS, 0, NULL, &event_NonMaxSup);
	errorController.CheckOpenCLError(clErr, "enquekernel nonMaxSupDDown");	
	
	clErr = clWaitForEvents(1, &event_NonMaxSup);
	errorController.CheckOpenCLError(clErr, "waitforevents nonMaxSupDDown");

// ========================================================================================================================================================================================
// ========================================================================================================================================================================================

	// zjisteni doby vypoctu
	clGetEventProfilingInfo(event_NonMaxSup, CL_PROFILING_COMMAND_START, sizeof(time_start), &time_start, NULL);
	clGetEventProfilingInfo(event_NonMaxSup, CL_PROFILING_COMMAND_END, sizeof(time_end), &time_end, NULL);
	total_time = (double)(time_end - time_start);

	// zapis zpracovaneho obrazu
	outImage = new cl_float[width * height];
	memset(outImage, 0, width * height * sizeof(cl_float));
	clErr = clEnqueueReadBuffer (*commandQueue, cl_outSup, CL_TRUE, 0, width * height * sizeof(cl_float), outImage, 0, NULL, NULL);
	errorController.CheckOpenCLError(clErr, "enqueread buffer NonMaxSup");

	// Uvolneni zdroju
	clErr = clReleaseMemObject(cl_in);
    errorController.CheckOpenCLError(clErr, "clReleaseMemObject input");    
    clErr = clReleaseMemObject(cl_outSup);
    errorController.CheckOpenCLError(clErr, "clReleaseMemObject output");
    clErr = clReleaseMemObject(cl_edgesDirection);
    errorController.CheckOpenCLError(clErr, "clReleaseMemObject edgesDirection");	

	// doba vypoctu v milisekundach
	return total_time / 1000000.0;
}